#ifndef __D_MINIMIZE1DCL_H__
#define __D_MINIMIZE1DCL_H__

#ifdef __D_VS_COMPILE__

#include <oclUtils.h>
#include "funcs.h"

cl_context				gpu_context;
cl_command_queue		command_que;
cl_program				program;

bool	is_init	= false;

void	init()
{
	gpu_context = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);

	size_t data_bytes_count	= 0;
	clGetContextInfo(gpu_context, CL_CONTEXT_DEVICES, 0, NULL, &data_bytes_count);
	cl_device_id* devices	= (cl_device_id*)malloc(data_bytes_count);
	clGetContextInfo(gpu_context, CL_CONTEXT_DEVICES, data_bytes_count, devices, NULL);
	command_que = clCreateCommandQueue(gpu_context, devices[0], 0, NULL);
	free(devices);

	size_t source_length	= 0;
	char* source_cl	= oclLoadProgSource("../GMcl/Common/minimize1dcl.h", "", &source_length);
	program = clCreateProgramWithSource(gpu_context, 1, (const char **)&source_cl, &source_length, NULL);
	free(source_cl);
	if (clBuildProgram(program, 0, NULL, NULL, NULL, NULL))
	{
		assert(0);
	}
}

void	minimize_fetch1dcl(test_func* _func, test2d* _fetch_array, args _args, size _argc, arg _min, arg _max, arg _eps)
{
	if (is_init == false)
	{
		init();
		is_init	= true;
	}

	cl_mem fetch_array_cl	= clCreateBuffer(gpu_context, CL_MEM_READ_WRITE, sizeof(test2d) * fetch_size2d, NULL, NULL);
	clEnqueueWriteBuffer(command_que, fetch_array_cl, CL_FALSE, 0, sizeof(test2d) * fetch_size2d, _fetch_array, 0, NULL, NULL);
	cl_mem args_cl			= clCreateBuffer(gpu_context, CL_MEM_READ_ONLY, sizeof(arg) * _argc, NULL, NULL);
	clEnqueueWriteBuffer(command_que, fetch_array_cl, CL_FALSE, 0, sizeof(arg) * _argc, _args, 0, NULL, NULL);

	cl_kernel kernel = clCreateKernel(program, "minimize", NULL);

	clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&fetch_array_cl);
	clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&args_cl);
	clSetKernelArg(kernel, 2, sizeof(cl_int), (void*)&_argc);
	clSetKernelArg(kernel, 3, sizeof(cl_float), (void*)&_min);
	clSetKernelArg(kernel, 4, sizeof(cl_float), (void*)&_max);
	clSetKernelArg(kernel, 5, sizeof(cl_float), (void*)&_eps);

	size_t	work_size	= fetch_size2d;
	clEnqueueNDRangeKernel(command_que, kernel, 1, NULL, &work_size, &work_size, 0, NULL, NULL);

	clEnqueueReadBuffer(command_que, fetch_array_cl, CL_TRUE, 0, sizeof(test2d) * fetch_size2d, _fetch_array, 0, NULL, NULL);
}

void	free()
{
}

#else

#define	value(array, i)		(array[3*i])
#define	coord(array, i, j)	(array[3*i + j + 1])

__kernel void	minimize(__global float* _fetch_array, __global const float* _args, 
						 int _argc, float _min, float _max, float _eps)
{
	int iGID = get_global_id(0);
	if (iGID >= 8)
	{   
		return;
	}

	value(_fetch_array, iGID)		= _eps;
	coord(_fetch_array, iGID, 0)	= _min;
	coord(_fetch_array, iGID, 1)	= _max;
}
#endif

#endif